You are an expert CUDA programmer tasked with accelerating a PyTorch model by replacing its operators with a highly optimized, custom CUDA kernel. You should consider operator fusion and algorithmic optimizations to achieve maximum speedup.

Here's an example to show you the syntax of inline embedding custom CUDA operators in PyTorch. The example given architecture is a simple ReLU:

python
import torch
import torch.nn as nn
import torch.nn.functional as F

class Model(nn.Module):
def init(self) -> None:
super().init()

def forward(self, x):
    return torch.relu(x)
def get_inputs():
x = torch.randn(1, 128).cuda()
return [x]

def get_init_inputs():
return []



The example new architecture with a custom CUDA kernel looks like this:

python
import torch
from torch.utils.cpp_extension import load_inline

relu_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

global void relu_kernel(const float* x, float* y, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
y[idx] = fmaxf(x[idx], 0.f);
}
}

torch::Tensor relu_cuda(torch::Tensor x) {
auto size = x.numel();
auto y = torch::empty_like(x);
const int block_size = 256;
int num_blocks = (size + block_size - 1) / block_size;
relu_kernel<<<num_blocks, block_size>>>(x.data_ptr<float>(), y.data_ptr<float>(), size);
return y;
}
"""

relu_cpp_source = """
torch::Tensor relu_cuda(torch::Tensor x);
"""

Compile the inline CUDA code
relu = load_inline(
name=“relu”,
cpp_sources=relu_cpp_source,
cuda_sources=relu_source,
functions=[“relu_cuda”],
verbose=True
)

class ModelNew(torch.nn.Module):
def init(self):
super(ModelNew, self).init()
self.relu = relu # The module containing the kernel

def forward(self, x):
    return self.relu.relu_cuda(x)
def get_inputs():
x = torch.randn(1, 128).cuda()
return [x]

def get_init_inputs():
return []



---

Now, you are given the following PyTorch architecture to accelerate. The model first applies Instance Normalization to a 4D input tensor and then computes the Minkowski norm of each channel in the normalized output. This baseline implementation uses standard PyTorch operations.

python
import torch
import torch.nn as nn

class Model(nn.Module):
“”"
Instance Normalization followed by a Minkowski Transform.
This version uses standard PyTorch operations for a fair baseline.
“”"
def init(self, num_features=64, p=2, eps=1e-5):
super(Model, self).init()
self.p = p
if p <= 0:
raise ValueError(“p must be positive”)

    # InstanceNorm2d operates on [N, C, H, W]
    self.instance_norm = nn.InstanceNorm2d(num_features, eps=eps, affine=False, track_running_stats=False)

def forward(self, x: torch.Tensor) -> torch.Tensor:
    """
    Compute the InstanceNorm of x, then compute the Minkowski transform of the result.
    Input: [N, C, H, W]
    Output: [N, C]
    """
    # Input validation
    if x.dim() != 4:
        raise ValueError(f"Input tensor must be 4D, got {x.dim()}D")
    
    # Step 1: Apply Instance Normalization
    normed_x = self.instance_norm(x) # Shape: [N, C, H, W]
    
    # Step 2: Compute Minkowski transform per channel on the normalized features
    N, C, H, W = normed_x.shape
    normed_x_flat = normed_x.view(N, C, H * W)
    abs_normed_flat = torch.abs(normed_x_flat)
    
    if self.p == 1:
        minkowski_vals = torch.sum(abs_normed_flat, dim=2) # Shape: [N, C]
    elif self.p == 2:
        minkowski_vals = torch.sqrt(torch.sum(abs_normed_flat ** 2, dim=2)) # Shape: [N, C]
    else:
        minkowski_vals = torch.pow(torch.sum(torch.pow(abs_normed_flat, self.p), dim=2), 1.0/self.p) # Shape: [N, C]
    
    return minkowski_vals
参数配置
batch_size = 32
num_features = 64
height = 128
width = 128

def get_inputs():
x = torch.randn(batch_size, num_features, height, width)
return [x]

def get_init_inputs():
return [num_features, 2] # num_features and p value

Your task is to generate the `ModelNew` architecture with a custom CUDA kernel that fuses the Instance Normalization calculation and the Minkowski norm calculation into a single kernel launch.

**CRITICAL REQUIREMENTS:**

1.  **Operator Fusion:** The entire logic—computing instance mean and variance, normalizing the features, and then calculating the Minkowski norm for each channel—must be performed inside a **single CUDA kernel**. No intermediate tensors (like the normalized feature map) should be written to global memory.
2.  **Kernel Logic:**
    *   Each thread block should be responsible for computing the final output for a single channel of a single sample.
    *   The kernel should use a `dim3` grid where `blockIdx.x` is the sample index (N) and `blockIdx.y` is the channel index (C).
    *   The kernel must perform three main stages:
        a. Calculate the `mean` and `inv_std` (1/sqrt(var + eps)) for the target channel. This requires two passes over the data (or a more complex single-pass algorithm) with efficient block-wide reductions using `extern __shared__`.
        b. Use the calculated `mean` and `inv_std` to normalize each element of the channel.
        c. Calculate the Minkowski norm of the normalized channel elements. This also requires a block-wide reduction.
3.  **Efficient Reduction:** You must implement an efficient block-wide reduction function (e.g., `blockReduceSum`) using shared memory to compute the sums required for mean, variance, and the Minkowski norm.
4.  **Final Calculation:** Inside the kernel, after the Minkowski norm is computed by the first thread (`tid == 0`), the final result should be written to the output tensor. The output tensor should have shape `[N, C]`.
5.  **Performance Optimization:** The host-side function should launch the kernel with an appropriate number of threads per block (e.g., 256) and allocate the required shared memory.
6.  **Code Structure:** Follow the exact structure of the provided example, including `load_inline`, the CUDA source string, the C++ wrapper source string, and the `ModelNew` class. The `get_init_inputs` function must return `[num_features, 2]` to match the baseline.
7.  **No Fast Math:** Do not use `--use_fast_math` in the compilation flags to ensure numerical accuracy.
